feat(xpu): add Intel GPU (XPU) support#1110
Open
aslanxie wants to merge 1 commit into
Open
Conversation
Contributor
There was a problem hiding this comment.
Pull request overview
Note
Copilot was unable to run its full agentic suite in this review.
This PR adds Intel XPU (PyTorch XPU) support across runtime device detection, Python model components, C++/pybind ops, and Bazel build/packaging so the project can run on Intel GPUs with appropriate fallbacks.
Changes:
- Introduce XPU device type and GPU-agnostic helpers (availability/count/device selection/visible devices).
- Add XPU-specific Python modules (attention SDPA + vLLM kernels wrapper, norms/activations/linear strategies) and KV-cache layout handling.
- Extend C++ runtime/ops and Bazel toolchain + wheel metadata to support
--config=xpubuilds (Python 3.12, SYCL toolchain, XPU bindings).
Reviewed changes
Copilot reviewed 92 out of 96 changed files in this pull request and generated no comments.
Show a summary per file
| File | Description |
|---|---|
| rtp_llm/start_backend_server.py | Switch GPU detection to device-agnostic helpers and add VIT separation server path. |
| rtp_llm/ops/init.py | Add XPU detection logging and make libpython preload version-agnostic. |
| rtp_llm/models_py/utils/arch.py | Extend device-type utilities import to include XPU helpers. |
| rtp_llm/models_py/standalone/auto_model.py | Select xpu device when available; adjust KV layout and pin_memory behavior for XPU. |
| rtp_llm/models_py/modules/hybrid/causal_attention.py | Add XPU norm import for hybrid causal attention. |
| rtp_llm/models_py/modules/factory/linear/impl/xpu/f16_linear.py | Add XPU F16/BF16 Linear backend using PyTorch F.linear. |
| rtp_llm/models_py/modules/factory/linear/impl/xpu/init.py | Register XPU Linear strategies in the factory. |
| rtp_llm/models_py/modules/factory/linear/init.py | Route Linear factory registration to XPU strategies when on XPU. |
| rtp_llm/models_py/modules/factory/fused_moe/impl/xpu/init.py | Add XPU MoE placeholder module. |
| rtp_llm/models_py/modules/factory/fused_moe/init.py | Configure MoE registry for XPU to use batched Triton fallback. |
| rtp_llm/models_py/modules/factory/attention/xpu_impl/test/test_kv_cache_layout.py | Add CPU-runnable test guarding XPU KV cache NSHD layout contract. |
| rtp_llm/models_py/modules/factory/attention/xpu_impl/test/BUILD | Bazel target for KV cache layout test. |
| rtp_llm/models_py/modules/factory/attention/xpu_impl/sdpa.py | Add XPU SDPA attention implementations for prefill/decode with RoPE + paged cache. |
| rtp_llm/models_py/modules/factory/attention/xpu_impl/init.py | Add XPU attention package marker. |
| rtp_llm/models_py/modules/factory/attention/init.py | Register XPU attention implementations in the attention factory lists. |
| rtp_llm/models_py/modules/base/xpu/vllm_xpu_ops.py | Add wrapper for optional vllm-xpu-kernels ops with PyTorch fallbacks. |
| rtp_llm/models_py/modules/base/xpu/not_implemented_ops.py | Add XPU stubs for unsupported ops. |
| rtp_llm/models_py/modules/base/xpu/norm.py | Add XPU norm implementations with optional vllm-xpu-kernels acceleration. |
| rtp_llm/models_py/modules/base/xpu/moe_gating.py | Add PyTorch fallback MoE gating op for XPU. |
| rtp_llm/models_py/modules/base/xpu/activation.py | Add XPU fused SiLU-and-mul implementation with optional kernel acceleration. |
| rtp_llm/models_py/modules/base/common/embedding.py | Add fallback path when compiled embedding op is unavailable. |
| rtp_llm/models_py/modules/base/init.py | Wire base module imports for XPU device type. |
| rtp_llm/models_py/bindings/xpu/XpuTorchExt.h | Add XPU-specific torch extension header. |
| rtp_llm/models_py/bindings/xpu/RegisterXpuOps.cc | Register XPU pybind ops entry point. |
| rtp_llm/models_py/bindings/xpu/RegisterXpuBaseBindings.hpp | Provide XPU fallback implementations for key kernels in bindings. |
| rtp_llm/models_py/bindings/xpu/BUILD | Bazel target building XPU bindings. |
| rtp_llm/models_py/bindings/core/ExecOps.h | Add getTorchDevice() API and keep CUDA alias for compatibility. |
| rtp_llm/models_py/bindings/core/ExecOps.cc | Extend runtime sync/event/device/memory queries to XPU. |
| rtp_llm/models_py/bindings/core/CudaSampleOp.cc | Add XPU pure-PyTorch sampling implementation and disable speculative sampling. |
| rtp_llm/models_py/bindings/core/CudaOps.cc | Add XPU implementations for copy and logits masking operations. |
| rtp_llm/models_py/bindings/core/CudaBeamSearchOp.cc | Add PyTorch fallback beam search for XPU. |
| rtp_llm/models_py/bindings/core/BUILD | Update core bindings build graph for XPU selects and SYCL feature flags. |
| rtp_llm/models_py/bindings/common/kernels/BUILD | Disable CUDA-only fuse-copy kernel compilation on XPU. |
| rtp_llm/models_py/bindings/common/FusedCopyOp.cc | Add XPU fallback for fused copy ops using SYCL queue memcpy. |
| rtp_llm/models_py/bindings/common/BUILD | Adjust common bindings build to include/exclude CUDA-only sources on XPU. |
| rtp_llm/models_py/bindings/OpDefs.h | Add XPU KV cache NSHD layout and add position_ids field to attention inputs. |
| rtp_llm/models_py/bindings/OpDefs.cc | Expose new position_ids binding and make decode_cu_seqlens_host read-only. |
| rtp_llm/models/base_model.py | Prefer XPU device string when available. |
| rtp_llm/model_loader/weight_manager.py | Disable CUDA stream usage on XPU and adjust synchronization paths. |
| rtp_llm/model_loader/loader.py | Extend memory cleanup helper to XPU. |
| rtp_llm/frontend/frontend_app.py | Add uvicorn import fallback for loop auto setup. |
| rtp_llm/device/device_type.py | Add XPU device type and detection helper is_xpu(). |
| rtp_llm/device/device_impl.py | Add XPU device implementation and GPU-agnostic helper APIs. |
| rtp_llm/device/init.py | Register XPU device class in device factory. |
| rtp_llm/cpp/utils/TensorDebugUtils.h | Treat XPU tensors like CUDA for debug-dump restrictions. |
| rtp_llm/cpp/utils/ErrorCode.h | Include <string> explicitly (XPU toolchain include differences). |
| rtp_llm/cpp/pybind/th_utils.h | Make CUDA-check macros accept XPU tensors when building for XPU. |
| rtp_llm/cpp/pybind/ComputeInit.cc | Enable exec ctx ops registration for XPU builds. |
| rtp_llm/cpp/pybind/BUILD | Link XPU exec ops and adjust deps for XPU builds. |
| rtp_llm/cpp/normal_engine/speculative/SpeculativeSampler.cc | Use getTorchDevice() and treat XPU like CUDA for host copies. |
| rtp_llm/cpp/normal_engine/speculative/MtpExecutor.cc | Use getTorchDevice() for speculative buffers on XPU-enabled builds. |
| rtp_llm/cpp/normal_engine/speculative/MtpBatchStreamProcessor.cc | Update device transfers and CPU staging for XPU tensors. |
| rtp_llm/cpp/normal_engine/NormalSamplerInputGatherer.cc | Allocate all_probs on getTorchDevice() (CUDA/XPU). |
| rtp_llm/cpp/normal_engine/NormalOutputDispatcher.cc | Move label tensor to getTorchDevice() for loss computation. |
| rtp_llm/cpp/normal_engine/NormalModelInputGatherer.cc | Use getTorchDevice() for multimodal tensors in context batching. |
| rtp_llm/cpp/normal_engine/NormalEngine.cc | Add XPU caching allocator sync/empty-cache and warmup gating. |
| rtp_llm/cpp/models/logits_processor/MultiSeqLogitsProcessor.cc | Move mask to getTorchDevice() for XPU compatibility. |
| rtp_llm/cpp/models/logits_processor/BaseLogitsProcessor.cc | Return vocab mask on getTorchDevice() for XPU compatibility. |
| rtp_llm/cpp/models/eplb/ExpertBalancer.cc | Allocate tensors on getTorchDevice() for XPU compatibility. |
| rtp_llm/cpp/models/Sampler.cc | Switch sampler tensors/transfers to getTorchDevice() and fix variable-beam token copy. |
| rtp_llm/cpp/models/PyWrappedModel.h | Disable CUDA graph/prefill-CP features on XPU; add device sync for XPU. |
| rtp_llm/cpp/models/PyWrappedModel.cc | Generalize host->device tensor staging to getTorchDevice() and treat XPU as device. |
| rtp_llm/cpp/models/ModelTypes.cc | Allocate packed GPU buffers on getTorchDevice() and treat XPU as device. |
| rtp_llm/cpp/models/BUILD | Adjust model library deps for XPU builds (no CUDA graph impl; keep copy op). |
| rtp_llm/cpp/engine_base/stream/GenerateStream.cc | Add XPU generator support; treat XPU like CUDA for CPU staging. |
| rtp_llm/cpp/engine_base/WeightsConverter.cc | Copy tensors to getTorchDevice() for XPU compatibility. |
| rtp_llm/cpp/engine_base/TorchProfiler.h | Enable XPU profiler activity type when building for XPU. |
| rtp_llm/cpp/cache/connector/p2p/transfer/tcp/CudaCopyUtil.cc | Use getTorchDevice() for wrapped raw pointers in copies. |
| rtp_llm/cpp/cache/connector/p2p/LayerBlockConverterImpl.h | Treat XPU like CUDA in BlockInfo device classification. |
| rtp_llm/cpp/cache/connector/memory/KVCacheMemoryConnector.cc | Use getTorchDevice() for mem/gpu block tensor wrappers. |
| rtp_llm/cpp/cache/MemoryLayoutStrategy.cc | Treat XPU device tensors as GPU blocks. |
| rtp_llm/cpp/cache/MemoryEvaluationHelper.cc | Add XPU free/total memory query path. |
| rtp_llm/cpp/cache/KVCacheManager.cc | Treat XPU tensors as device sources/dests in KV updates. |
| rtp_llm/cpp/cache/BlockPool.cc | Allocate device-side block pool on getTorchDevice(); treat XPU as GPU. |
| rtp_llm/config/server_config_setup.py | Extend local world size/device setup to XPU and add fail-fast for XPU speculative decoding. |
| rtp_llm/BUILD | Add XPU-aware wheel requirements filtering and cp312 wheel tag target. |
| deps/requirements_xpu.txt | Add standalone requirements list for XPU environment (Python 3.12, XPU torch index). |
| deps/pip.bzl | Add pip_parse for XPU lockfile and XPU extra-index URL. |
| deps/BUILD | Add target to compile XPU lockfile. |
| bazel/device_defs.bzl | Add XPU test env selection. |
| bazel/defs.bzl | Allow wheel renaming with configurable Python tag (cp312 for XPU). |
| arch_config/arch_select.bzl | Add XPU dependency selection, wheel req filtering/remap/overrides, and torch deps for XPU. |
| WORKSPACE | Add XPU configure rules and torch_xpu repository; load XPU pip deps. |
| BUILD.pytorch | Add using_xpu config and link XPU runtime libraries + python headers for XPU. |
| BUILD | Add using_xpu config_setting. |
| 3rdparty/gpus/xpu_python_utils.bzl | Add helper to resolve symlinked python inside venvs for repo rules. |
| 3rdparty/gpus/xpu_configure.bzl | Add Intel oneAPI/SYCL toolchain auto-configuration and Python 3.12 validation for XPU builds. |
| 3rdparty/gpus/xpu/BUILD.tpl | Add template build targets for SYCL runtime + Level Zero loader. |
| 3rdparty/gpus/torch_xpu_configure.bzl | Add repository rule to locate system-installed PyTorch XPU site-packages. |
| 3rdparty/gpus/crosstool/xpu_cc_toolchain_config.bzl.tpl | Add cc_toolchain_config for SYCL compilation/linking flags. |
| 3rdparty/gpus/crosstool/clang/bin/crosstool_wrapper_driver_xpu.tpl | Add crosstool wrapper routing Bazel C/C++ to icx/icpx with flag filtering. |
| .bazelrc | Add --config=xpu build/test settings for SYCL toolchain, env vars, and Python path. |
Comments suppressed due to low confidence (7)
rtp_llm/models_py/bindings/core/CudaSampleOp.cc:1
- XPU sampling reinterprets
top_k(int32) asuint32_t, which breaks semantics for disabled values (e.g.,top_k <= 0). A negativetop_kbecomes a hugeuint32_t, causinghas_top_k/kcomputation to behave incorrectly and potentially calltopk()with unintendedk. Use anint32_t*(orint64_t) view fortop_kchecks/clamping, and avoidreinterpret_cast<uint32_t*>here.
rtp_llm/models_py/bindings/core/CudaSampleOp.cc:1 - XPU sampling reinterprets
top_k(int32) asuint32_t, which breaks semantics for disabled values (e.g.,top_k <= 0). A negativetop_kbecomes a hugeuint32_t, causinghas_top_k/kcomputation to behave incorrectly and potentially calltopk()with unintendedk. Use anint32_t*(orint64_t) view fortop_kchecks/clamping, and avoidreinterpret_cast<uint32_t*>here.
rtp_llm/models_py/bindings/core/CudaSampleOp.cc:1 - XPU sampling reinterprets
top_k(int32) asuint32_t, which breaks semantics for disabled values (e.g.,top_k <= 0). A negativetop_kbecomes a hugeuint32_t, causinghas_top_k/kcomputation to behave incorrectly and potentially calltopk()with unintendedk. Use anint32_t*(orint64_t) view fortop_kchecks/clamping, and avoidreinterpret_cast<uint32_t*>here.
rtp_llm/models_py/modules/factory/attention/xpu_impl/sdpa.py:1 - There are repeated
.cpu()conversions inside request loops (block_ids_all[req_idx].cpu(),block_ids_all[0].cpu(),block_ids_all[i].cpu()), which can introduce per-iteration overhead and synchronization. Moveblock_ids_allto a CPU tensor once (if needed) before the loop, then index it without further device transfers; likewise, only computebidson CPU once per forward path.
rtp_llm/models_py/modules/factory/attention/xpu_impl/sdpa.py:1 - There are repeated
.cpu()conversions inside request loops (block_ids_all[req_idx].cpu(),block_ids_all[0].cpu(),block_ids_all[i].cpu()), which can introduce per-iteration overhead and synchronization. Moveblock_ids_allto a CPU tensor once (if needed) before the loop, then index it without further device transfers; likewise, only computebidson CPU once per forward path.
rtp_llm/models_py/modules/base/xpu/vllm_xpu_ops.py:1 - Inserting an arbitrary environment-controlled path at the front of
sys.pathcan enable unintended module shadowing/import hijacking. Prefer loading the extension via a controlled mechanism (e.g., validating the path is absolute/expected, warning when enabled, or usingimportlibwith a targeted loader) rather than globally modifying import precedence.
rtp_llm/models_py/modules/base/common/embedding.py:1 - When the compiled
rtp_llm_ops.embeddingis unavailable, the fallback path silently ignorestext_tokens_mask(multimodal masking) and proceeds, which can produce incorrect model outputs. A warning-once is easy to miss in production; consider failing fast whentext_tokens_maskis provided (or implementing mask support in the fallback) to avoid silently wrong results.
💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.
Collaborator
AI Code Review - PR #1110Status: BLOCKING Summary: P0/0 · P1/7 · P2/15 · P3/0 Blocking IssuesP1
Non-blocking SuggestionsP2
Checklist ✅ (56 items passed)Strengths
|
aslanxie
force-pushed
the
feat/xpu-support
branch
from
June 17, 2026 03:03
3060828 to
e9f12a7
Compare
Collaborator
AI Code Review - PR #1110Status: BLOCKING Summary: P0/0 · P1/3 · P2/14 · P3/0 Blocking IssuesP1
Non-blocking SuggestionsP2
Checklist ✅ (56 items passed)Strengths
|
aslanxie
force-pushed
the
feat/xpu-support
branch
from
June 17, 2026 08:07
e9f12a7 to
1ba8d0a
Compare
Contributor
There was a problem hiding this comment.
Pull request overview
Copilot reviewed 93 out of 97 changed files in this pull request and generated 1 comment.
Comments suppressed due to low confidence (5)
rtp_llm/device/device_impl.py:1
_is_xpu_device()ignores the newRTP_LLM_DEVICE_TYPEoverride logic inget_device_type()and will return true whenevertorch.xpu.is_available()is true, even if the user explicitly forced CUDA. This can route server setup and device selection down the XPU path unexpectedly. Prefer basing this onget_device_type() == DeviceType.Xpu/is_xpu()(or passing an already-resolved DeviceType) so override + detection are consistent everywhere.
rtp_llm/models_py/modules/base/common/embedding.py:1- This fallback silently ignores
text_tokens_mask(only a one-time warning), which can produce incorrect multimodal outputs. This is also inconsistent with the XPU C++ binding added inRegisterXpuBaseBindings.hpp, which hard-fails whentext_tokens_maskis provided. To avoid silent correctness issues, make the Python fallback reject non-emptytext_tokens_mask(raise), or implement equivalent masking behavior in the fallback so semantics match the fused op.
rtp_llm/models_py/modules/factory/attention/init.py:1 - This change drops the previous ordering logic that explicitly kept
XQAImplhigher-priority to avoid token divergence and golden refreshes (per the removed comment). Ifget_xqa_impl()returns a different implementation, decode behavior and numerics can change compared to prior releases. Consider restoring the old behavior: appendXQAImplfirst, then appendget_xqa_impl()only when it differs, so the default remains stable unless explicitly changed.
rtp_llm/frontend/frontend_app.py:1 auto_loop_factoryis not a drop-in replacement forauto_loop_setupin uvicorn; it typically returns a loop implementation rather than performing setup. If the rest of this module expectsauto_loop_setup(...)side effects, this fallback can break event-loop initialization at runtime. Prefer defining a small compatibility wrapper that preserves the expected call semantics (e.g., call the factory and then apply the result to the uvicorn config), or gate on uvicorn version with the correct API for each.
rtp_llm/models_py/bindings/core/CudaSampleOp.cc:1- Reinterpreting an
int32_t*buffer asuint32_t*can violate C++ strict-aliasing rules and is undefined behavior in optimized builds. Use the nativeint32_t*pointer and assign1directly (or perform a safe cast per element) to avoid UB while keeping the same comparison/branch behavior.
Comment on lines
+97
to
+100
| elif normalized in _XPU_PACKAGE_REMAP: | ||
| xpu_reqs.append(_XPU_PACKAGE_REMAP[normalized]) | ||
| else: | ||
| xpu_reqs.append(req) |
aslanxie
force-pushed
the
feat/xpu-support
branch
from
June 17, 2026 08:46
1ba8d0a to
32d1d4e
Compare
Collaborator
AI Code Review - PR #1110Status: BLOCKING Summary: P0/0 · P1/2 · P2/21 · P3/0 Blocking IssuesP1
Non-blocking SuggestionsP2
Checklist ✅ (56 items passed)Strengths
|
aslanxie
force-pushed
the
feat/xpu-support
branch
from
June 17, 2026 14:54
32d1d4e to
f79e7cb
Compare
Contributor
There was a problem hiding this comment.
Pull request overview
Copilot reviewed 93 out of 97 changed files in this pull request and generated no new comments.
Comments suppressed due to low confidence (4)
rtp_llm/device/device_impl.py:1
gpu_is_available()currently returns true based on the selected device type (includingRTP_LLM_DEVICE_TYPEoverride), not on actual runtime availability. If a user forcesRTP_LLM_DEVICE_TYPE=xpuon a build withouttorch.xpu(or forcescudawhen CUDA isn’t available),gpu_device_count()will raise (torch.xpumissing) or return 0 and downstream code can hit division-by-zero / invalid world-size checks. Fix by gating onhasattr(torch, 'xpu') and torch.xpu.is_available()/torch.cuda.is_available()insidegpu_is_available()andgpu_device_count()(and ideally raise a clear error when an override requests an unavailable backend).
rtp_llm/device/device_impl.py:1- Parsing
ZE_AFFINITY_MASKentries withint(visible[local_rank])will fail for valid Level Zero affinity formats like0.0/0.1(device.tile). IfZE_AFFINITY_MASKcontains tile-qualified entries, this code will throw and prevent startup. Consider parsing by splitting on'.'(taking the device portion) or otherwise supporting tile notation explicitly, and document the expected format.
rtp_llm/models_py/modules/base/common/embedding.py:1 - When
text_tokens_maskis provided, silently ignoring it produces incorrect embeddings for multimodal masked inputs. A warning is easy to miss and turns a correctness requirement into best-effort behavior. Prefer raising a clear exception whentext_tokens_maskis non-empty and the fusedrtp_llm_ops.embeddingop is unavailable, so masked multimodal runs fail fast instead of returning wrong outputs.
rtp_llm/start_backend_server.py:1 _get_cuda_device_list()now returns a generic GPU/XPU-visible list (viaget_visible_device_list()), so the function name is misleading and increases confusion in XPU paths (especially where it later feedsZE_AFFINITY_MASK). Renaming it (e.g.,_get_visible_gpu_device_list) and updating the corresponding local variable names (e.g.,cuda_device_list) would reduce backend-specific ambiguity.
Collaborator
AI Code Review - PR #1110Status: BLOCKING Summary: P0/0 · P1/7 · P2/14 · P3/0 Blocking IssuesP1
Non-blocking SuggestionsP2
Checklist ✅ (56 items passed)Strengths
|
aslanxie
added a commit
to aslanxie/rtp-llm-xpu
that referenced
this pull request
Jun 17, 2026
- embedding.py: fail-fast when text_tokens_mask/position_ids/token_types present but rtp_llm_ops.embedding unavailable (no silent wrong output) - server_config_setup.py: route set_device through gpu_set_device() to honor RTP_LLM_DEVICE_TYPE override on mixed XPU+CUDA hosts - base_model.py: use get_device_string() for resolved device type - weight_manager.py: use _is_xpu_device() instead of raw torch.xpu.is_available() for stream/sync selection - start_backend_server.py: hoist XPU world_size>1 fail-fast above device_count branch (covers single-card case); add gpu_device_count()==0 guard against div-by-zero on invalid RTP_LLM_DEVICE_TYPE override - vllm_flash_attn.py: replace weak sum+last prefill cache key with full content digest; add block-table fingerprint to decode cache keys (_write_idx, _flat_bids, _seqused_k) for hybrid-model safety Addresses P1 items #2-#7 from LLLLKKKK's AI code review. Item #1 (per-layer KV gather) is a tracked follow-up: XPU FA2 kernel requires contiguous pages and the interleaved PD layout prevents direct consumption.
aslanxie
added a commit
to aslanxie/rtp-llm-xpu
that referenced
this pull request
Jun 17, 2026
- embedding.py: fail-fast when text_tokens_mask/position_ids/token_types present but rtp_llm_ops.embedding unavailable (no silent wrong output) - server_config_setup.py: route set_device through gpu_set_device() to honor RTP_LLM_DEVICE_TYPE override on mixed XPU+CUDA hosts - base_model.py: use get_device_string() for resolved device type - weight_manager.py: use _is_xpu_device() instead of raw torch.xpu.is_available() for stream/sync selection - start_backend_server.py: hoist XPU world_size>1 fail-fast above device_count branch (covers single-card case); add gpu_device_count()==0 guard against div-by-zero on invalid RTP_LLM_DEVICE_TYPE override - vllm_flash_attn.py: replace weak sum+last prefill cache key with full content digest; add block-table fingerprint to decode cache keys (_write_idx, _flat_bids, _seqused_k) for hybrid-model safety Addresses P1 items #2-#7 from LLLLKKKK's AI code review. Item #1 (per-layer KV gather) is a tracked follow-up: XPU FA2 kernel requires contiguous pages and the interleaved PD layout prevents direct consumption.
Collaborator
AI Code Review - PR #1110Status: BLOCKING Summary: P0/0 · P1/2 · P2/18 · P3/1 Blocking IssuesP1
Non-blocking SuggestionsP2
P3
Checklist Violations (6 fail / 56 total)General Principles Checklist
Python Static-First Checklist
Strengths
|
Contributor
There was a problem hiding this comment.
Pull request overview
Copilot reviewed 93 out of 97 changed files in this pull request and generated no new comments.
Comments suppressed due to low confidence (7)
rtp_llm/models_py/bindings/core/CudaOps.cc:1
c10::xpu::getCurrentXPUStream()is typically an XPUStream object, not asycl::queue&. Treating it as asycl::queuewill likely fail to compile on XPU builds. Use the correct API to obtain the underlying SYCL queue (e.g., get a stream object first and then call.queue()/ equivalent) or use a PyTorch-provided copy primitive for device-to-device copies.
rtp_llm/models_py/bindings/common/FusedCopyOp.cc:1- Same issue as in
CudaOps.cc:getCurrentXPUStream()is unlikely to be asycl::queue&, so this XPU fallback will likely not compile. Retrieve the underlying queue from the XPU stream via the correct accessor, or replace this path withat::Tensor/copy_-based copying that doesn't depend on SYCL types.
rtp_llm/models_py/bindings/core/CudaBeamSearchOp.cc:1 Tensor.where(...)is Python-style and is not a stable/portable libtorch C++ API call; on many PyTorch versions this will not compile (C++ typically usestorch::where(condition, a, b)). Replace with the C++ API form to avoid XPU build failures.
rtp_llm/device/device_impl.py:1gpu_is_available()currently returnsTruebased on the resolved device type (including viaRTP_LLM_DEVICE_TYPEoverride), even if that backend is not actually usable (e.g., override set tocudaon a host withtorch.cuda.device_count()==0). This can send the code down GPU-only paths and fail later. Prefer defining availability in terms ofgpu_device_count() > 0(or validate the override maps to an available backend).
rtp_llm/models/base_model.py:1- This returns
cpu:<rank>when running on CPU, which is not a standard device string for PyTorch tensor placement (CPU typically uses justcpu). Consider special-casing CPU to returncpu(no ordinal) while keepingcuda:<rank>/xpu:<rank>for GPU backends.
rtp_llm/models_py/modules/factory/attention/xpu_impl/sdpa.py:1 - This prefill path performs per-request
.cpu()transfers inside a Python loop (bids = block_ids_all[req_idx].cpu()), which can introduce repeated device↔host sync ifblock_ids_allis device-resident. Since you already prefer host block-id tensors (and SDPA code aims to avoid D2H sync), consider normalizingblock_ids_allto a CPU tensor exactly once (or using the existing host copies consistently) before entering the loop.
rtp_llm/start_backend_server.py:1 _get_cuda_device_list()now returns a device-agnostic GPU visible list (CUDA/ROCm/XPU). The function name is misleading after this change. Renaming it to something like_get_visible_gpu_device_list()(and updating call sites within this file) would make the intent clearer and reduce confusion.
Collaborator
AI Code Review - PR #1110Status: BLOCKING Summary: P0/0 · P1/4 · P2/15 · P3/0 Blocking IssuesP1
Non-blocking SuggestionsP2
Checklist ✅ (56 items passed)Strengths
|
Collaborator
AI Code Review - PR #1110Status: BLOCKING Summary: P0/0 · P1/5 · P2/20 · P3/0 Blocking IssuesP1
Non-blocking SuggestionsP2
Checklist ✅ (56 items passed)Strengths
|
Contributor
There was a problem hiding this comment.
Pull request overview
Copilot reviewed 94 out of 98 changed files in this pull request and generated no new comments.
Comments suppressed due to low confidence (5)
rtp_llm/start_backend_server.py:1
_get_cuda_device_listno longer returns CUDA-only devices (it now returns a device-type-dependent visible list). This name is misleading and makes call sites harder to reason about; rename it (and related variables likecuda_device_list) to something device-agnostic such as_get_gpu_device_list/gpu_device_list.
rtp_llm/models_py/modules/factory/attention/xpu_impl/sdpa.py:1torch.Tensordoes not consistently expose anis_cpuattribute across PyTorch versions/configurations (whereasis_cudais standard). To avoid attribute errors at runtime, usecu_seqlens.device.type == 'cpu'(or equivalent device-type checks) instead ofcu_seqlens.is_cpuin this module (same applies to otheris_cpuuses here).
rtp_llm/models_py/bindings/core/CudaOps.cc:1c10::xpu::getCurrentXPUStream()is used elsewhere in this PR as a stream object with.synchronize(), but here it's treated as asycl::queue&. This inconsistency is likely a compile-time type error (or at minimum relies on a non-obvious implicit conversion). Prefer extracting the underlying SYCL queue explicitly from the XPU stream (per the API provided by the XPU stream type) and use that formemcpy.
rtp_llm/models_py/bindings/common/FusedCopyOp.cc:1- Same issue as in
CudaOps.cc:c10::xpu::getCurrentXPUStream()is treated as asycl::queue&, which likely does not match the actual return type (and is inconsistent with.synchronize()usage elsewhere). Use the proper API to obtain the underlying SYCL queue from the XPU stream before callingmemcpy.
rtp_llm/models_py/modules/factory/attention/xpu_impl/sdpa.py:1 - The new XPU SDPA impl introduces non-trivial gating and behavioral constraints (e.g., rejecting prefix-cache hits, rejecting non-BASE KV cache dtype, rejecting unsupported RoPE styles). The PR adds XPU attention helper tests, but there are no unit tests shown covering these
support()decisions. Add targeted tests to lock in the selection behavior (especially the prefix_lengths>0 rejection) to prevent factory regressions.
Collaborator
AI Code Review - PR #1110Status: BLOCKING Summary: P0/0 · P1/4 · P2/22 · P3/0 Blocking IssuesP1
Non-blocking SuggestionsP2
Checklist ✅ (56 items passed)Strengths
|
Collaborator
AI Code Review - PR #1110Status: BLOCKING Summary: P0/0 · P1/3 · P2/15 · P3/0 Blocking IssuesP1
Non-blocking SuggestionsP2
Checklist ✅ (56 items passed)Strengths
|
Comment on lines
+645
to
+649
| auto selected = torch::multinomial(filtered_probs, 1, false).squeeze(-1); | ||
| // For any degenerate row, override the (uniform) draw with argmax on the | ||
| // original logits. Pure-device select, no host sync. | ||
| auto fallback = torch::argmax(params.logits, -1, false); | ||
| selected = torch::where(row_valid, selected, fallback); |
Collaborator
AI Code Review - PR #1110Status: BLOCKING Summary: P0/0 · P1/2 · P2/22 · P3/12 Blocking IssuesP1
Non-blocking SuggestionsP2
P3
Checklist Violations (2 fail / 56 total)General Principles Checklist
Python Static-First Checklist
Strengths
|
Collaborator
AI Code Review - PR #1110Status: LGTM Summary: P0/0 · P1/0 · P2/22 · P3/14 lgtm ready to ci Non-blocking SuggestionsP2
P3
Checklist ✅ (56 items passed)Strengths
|
Collaborator
|
internal source has been updated, please review the changes! |
Comment on lines
+434
to
+440
| if (params.do_sample.has_value()) { | ||
| auto top_k_ptr = params.top_k.data_ptr<int32_t>(); | ||
| bool any_greedy = false; | ||
| for (int64_t b = 0; b < batch_size; b++) { | ||
| if (!params.do_sample.value().data_ptr<bool>()[b]) { | ||
| top_k_ptr[b] = 1; | ||
| any_greedy = true; |
Comment on lines
+591
to
+595
| bool has_top_k = !std::all_of(top_k_ptr, top_k_ptr + batch_size, [](uint32_t t) { return t <= 0; }); | ||
| if (has_top_k) { | ||
| for (int64_t b = 0; b < batch_size; b++) { | ||
| int64_t k = top_k_ptr[b] <= 0 ? vocab_size_padded : (int64_t)top_k_ptr[b]; | ||
| // Clamp k into [1, vocab_size_padded] so an out-of-range top_k can |
Comment on lines
+582
to
+584
| // 4. Softmax -> probabilities | ||
| auto probs_t = torch::softmax(params.logits, -1); | ||
| params.logits.copy_(probs_t); |
Collaborator
AI Code Review - PR #1110Status: BLOCKING Summary: P0/0 · P1/1 · P2/44 · P3/14 Blocking IssuesP1
Non-blocking SuggestionsP2
虽然框架通常提供 2D block table（[num_reqs, max_blocks]），但批量路径已经显式处理 1D 情况，单请求路径也应该。当前 _write_to_paged_cache 的 n_valid != N 检查能防止静默错误，但报错信息不直观。
P3
Checklist ✅ (56 items passed)Strengths
|
Collaborator
AI Code Review - PR #1110Status: BLOCKING Summary: P0/0 · P1/1 · P2/36 · P3/16 Blocking IssuesP1
Non-blocking SuggestionsP2
P3
Checklist ✅ (56 items passed)Strengths
|
aslanxie
force-pushed
the
feat/xpu-support
branch
from
June 28, 2026 08:28
ecc47fc to
4ae3419
Compare
Comment on lines
104
to
108
| # frontend cannot load libpython3.10.so, so we need to load it manually | ||
| # In XPU env, we keep libpython3.10.so hardcoded and explicitly link it | ||
| # to the actual Python runtime library. | ||
| import sysconfig | ||
| from ctypes import cdll |
| auto safe_mask = (write_pos < max_seq_len).squeeze(-1); // [batch, beam_out] | ||
| write_pos = write_pos.clamp(0, max_seq_len - 1); | ||
| // Write tokens only for beams that have not overflowed (device-side where). | ||
| auto safe_ids = output_ids.where(safe_mask, token_ids_out.gather(2, write_pos).squeeze(-1)); |
Comment on lines
+176
to
+179
| if _is_xpu_device(): | ||
| os.environ["ZE_AFFINITY_MASK"] = ",".join(cuda_device_list) | ||
| else: | ||
| os.environ["CUDA_VISIBLE_DEVICES"] = ",".join(cuda_device_list) |
Collaborator
AI Code Review - PR #1110Status: LGTM Summary: P0/0 · P1/0 · P2/29 · P3/18 lgtm ready to ci Non-blocking SuggestionsP2
P3
Checklist ✅ (56 items passed)Strengths
|
Collaborator
|
internal source has been updated, please review the changes! |
Collaborator
AI Code Review - PR #1110Status: BLOCKING Summary: P0/0 · P1/3 · P2/39 · P3/12 Blocking IssuesP1
Non-blocking SuggestionsP2
P3
Checklist ✅ (56 items passed)Strengths
|
aslanxie
force-pushed
the
feat/xpu-support
branch
from
June 29, 2026 08:47
27d8878 to
d89ad3a
Compare
Comment on lines
+654
to
+656
| bool has_any_generator = std::any_of( | ||
| params.generator.begin(), params.generator.end(), | ||
| [](const c10::optional<at::Generator>& g) { return g.has_value() && g->defined(); }); |
Comment on lines
+229
to
+232
| write_pos = write_pos.clamp(0, max_seq_len - 1); | ||
| // Write tokens only for beams that have not overflowed (device-side where). | ||
| auto safe_ids = output_ids.where(safe_mask, token_ids_out.gather(2, write_pos).squeeze(-1)); | ||
| token_ids_out.scatter_(2, write_pos, safe_ids.unsqueeze(-1)); |
Collaborator
AI Code Review - PR #1110Status: BLOCKING Summary: P0/0 · P1/1 · P2/30 · P3/14 Blocking IssuesP1
Non-blocking SuggestionsP2
P3
Checklist ✅ (56 items passed)Strengths
|
aslanxie
force-pushed
the
feat/xpu-support
branch
from
June 30, 2026 03:14
8ee282a to
7757b2d
Compare
Collaborator
AI Code Review - PR #1110Status: BLOCKING Summary: P0/0 · P1/2 · P2/7 · P3/1 Blocking IssuesP1
Non-blocking SuggestionsP2
P3
Checklist Violations (5 fail / 56 total)General Principles Checklist
RTP-LLM Checklist
Strengths
|
Contributor
There was a problem hiding this comment.
Pull request overview
Copilot reviewed 95 out of 99 changed files in this pull request and generated 2 comments.
Comments suppressed due to low confidence (1)
rtp_llm/ops/init.py:110
- The manual
cdll.LoadLibrary(.../libpython3.10.so)hard-codes the Python minor version. The PR introduces an XPU build that targets Python 3.12 (cp312 wheel tag / XPU toolchain), so this will likely fail at import time unless the environment provides a compatibility symlink.
Load the runtime’s actual LDLIBRARY from sysconfig instead so the same code works for both 3.10 and 3.12 builds.
Comment on lines
142
to
145
| def _get_cuda_device_list() -> List[str]: | ||
| """Get CUDA device list from environment or hardware detection""" | ||
| cuda_devices = os.environ.get("CUDA_VISIBLE_DEVICES", None) | ||
| return ( | ||
| cuda_devices.split(",") | ||
| if cuda_devices is not None | ||
| else [str(i) for i in range(torch.cuda.device_count())] | ||
| ) | ||
| """Get GPU device list from environment or hardware detection""" | ||
| return get_visible_device_list() | ||
|
|
Comment on lines
+633
to
+637
| auto row_sums = filtered_probs.sum(-1, true); | ||
| // Guard against invalid probability distributions (all-zero / NaN / Inf rows). | ||
| // Fall back to argmax on the original logits for any degenerate row. | ||
| auto row_valid = (row_sums.squeeze(-1) > 0) & row_sums.squeeze(-1).isfinite(); | ||
| filtered_probs = filtered_probs / row_sums.clamp_min(1e-10f); |
Collaborator
AI Code Review - PR #1110Status: BLOCKING Summary: P0/0 · P1/1 · P2/1 · P3/1 Blocking IssuesP1
Non-blocking SuggestionsP2
P3
Checklist Violations (2 fail / 56 total)General Principles Checklist
Strengths
|
aslanxie
force-pushed
the
feat/xpu-support
branch
from
July 6, 2026 01:33
2e29105 to
bcbb080
Compare
- Add --config=xpu to .bazelrc; new xpu_configure.bzl (479 lines) that auto-detects the Intel oneAPI/SYCL toolchain and emits a Bazel C++ toolchain (xpu_cc_toolchain_config.bzl.tpl, crosstool_wrapper_driver_xpu.tpl) - Add torch_xpu_configure.bzl to locate PyTorch XPU headers/libs - Extend arch_select.bzl to select XPU platform; update BUILD, BUILD.pytorch, bazel/defs.bzl, bazel/device_defs.bzl with USING_XPU compile-time flag - Add pip_xpu_torch Bazel repo rule for XPU Python deps (requirements_xpu.txt, requirements_lock_xpu.txt); gate the rule behind TF_NEED_XPU so that `bazel sync` on non-XPU Python 3.10 containers never resolves XPU-only pins (e.g. scikit-learn==1.8.0 requires Python>=3.11) - NormalEngine: include XPUCachingAllocator; add DeviceGuard for ctor thread; enable prefill warm-up and emptyCache() via XPU caching allocator - PyWrappedModel: replace all hard-coded torch::kCUDA / .cuda() calls with getTorchDevice(); add kPinHostMem=false on XPU (pinned memory unsupported); guard kv_cache_block_id_host and block_id pin_memory() calls accordingly - KVCacheManager, BlockPool, MemoryEvaluationHelper: add XPU memory evaluation and KV cache layout support - Sampler, ModelTypes, ComputeInit: add XPU device branching - New bindings/xpu/: RegisterXpuBaseBindings.hpp registers ~40 ops as pure-PyTorch fallbacks on XPU (RMSNorm, LayerNorm, attention, sampling, KV cache ops, write_cache_store) - New modules/base/xpu/: activation.py, norm.py (RMSNorm/LayerNorm), moe_gating.py, select_topk.py, vllm_xpu_ops.py (flash_attn_varlen wrapper with SDPA fallback when vllm-xpu-kernels FA2 is absent) - New modules/factory/attention/xpu_impl/vllm_flash_attn.py (906 lines): full paged-KV-cache attention for prefill and decode using flash_attn_varlen; rejects unsupported RoPE styles (Yarn/Llama3/Mrope/Glm2/DynamicNTK); hoist decode hash computation to step boundary (layer_idx wrap-around detection) instead of recomputing per layer - New modules/factory/linear/impl/xpu/f16_linear.py and modules/factory/fused_moe/impl/xpu/ stubs - Add unit tests: test_kv_cache_layout.py, test_no_rope.py - Add XpuImpl class extending GpuImpl: get_device_id() with fallback to LOCAL_RANK + ZE_AFFINITY_MASK; _get_mem_info() via torch.xpu.mem_get_info; guard empty ZE_AFFINITY_MASK to avoid spurious device-0 targeting - device_type.py: add XPU enum value and auto-detection logic - server_config_setup.py: device selection via ZE_AFFINITY_MASK - start_backend_server.py: XPU startup path - model_loader/loader.py, weight_manager.py: XPU weight loading; use gpu_is_available() + XPU memory APIs for fastsafetensor progress log - frontend_app.py, base_model.py: XPU path plumbing
Collaborator
AI Code Review - PR #1110Status: BLOCKING Summary: P0/0 · P1/0 · P2/1 · P3/1 Non-blocking SuggestionsP2
P3
Checklist Violations (1 fail / 56 total)General Principles Checklist
Strengths
|
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Overview
Add Intel GPU (XPU) inference support to RTP-LLM, reusing vllm-xpu-kernels to optimize performance on Intel GPU.
The base environment is the intel/vllm Docker image.
Guiding principles:
--config=xpu— DO NOT break existing code logicChanges
1. Build Infrastructure
xpu_configure.bzl(analogous tocuda_configure).bazelrc --config=xpupreset with oneAPI compiler flags2. C++ Device Generalization
select()branches in BUILD filesxpu_sycl_compilefeature flag3. Python Device & Attention
4. Module Factories & Server Integration
Test Environment
How to Build